[mlir][GPU] Do not strip location info when lowering to NVVM #120432
Conversation
|
@llvm/pr-subscribers-mlir @llvm/pr-subscribers-mlir-gpu Author: Matthias Springer (matthias-springer) ChangesThis is needed for a subsequent commit that reads location information when lowering Full diff: https://github.com/llvm/llvm-project/pull/120432.diff 1 Files Affected:
diff --git a/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp b/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
index fb440756e0c1d5..20d7372eef85d5 100644
--- a/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
+++ b/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
@@ -73,7 +73,6 @@ void buildCommonPassPipeline(
//===----------------------------------------------------------------------===//
void buildGpuPassPipeline(OpPassManager &pm,
const mlir::gpu::GPUToNVVMPipelineOptions &options) {
- pm.addNestedPass<gpu::GPUModuleOp>(createStripDebugInfoPass());
ConvertGpuOpsToNVVMOpsOptions opt;
opt.useBarePtrCallConv = options.kernelUseBarePtrCallConv;
opt.indexBitwidth = options.indexBitWidth;
|
matthias-springer
changed the title
| //===----------------------------------------------------------------------===// | ||
| void buildGpuPassPipeline(OpPassManager &pm, | ||
| const mlir::gpu::GPUToNVVMPipelineOptions &options) { | ||
| pm.addNestedPass<gpu::GPUModuleOp>(createStripDebugInfoPass()); |
There was a problem hiding this comment.
Why was this added in the first place?
There was a problem hiding this comment.
Not sure. Looking at the git history (https://reviews.llvm.org/D155463), there was a test case that had it in the pass pipeline. See also https://reviews.llvm.org/D103187. Maybe @ThomasRaoux knows.
There was a problem hiding this comment.
I remember that passing ptx with debug info used to crash nvidia driver. I had not debugged it at the time though
There was a problem hiding this comment.
Good to know. Given that the printf integration test works on my machine, this seems to be resolved. Let's see if the build bot complains after I merge this...
matthias-springer
deleted the
users/matthias-springer/do_not_strip_debug
branch
This is needed for a subsequent commit that reads location information when lowering
gpu.assert.